#include "data_utils.h"
#ifndef ASCENDC_CPU_DEBUG
#include "acl/acl.h"
extern void masked_fill_custom_do(uint32_t blockDim, 
                                  void *stream, 
                                  uint8_t *x, 
                                  uint8_t *mask, 
                                  uint8_t *value, 
                                  uint8_t *z);
#else
#include "tikicpulib.h"
extern "C" __global__ __aicore__ void masked_fill_custom(GM_ADDR x, 
                                                         GM_ADDR mask, 
                                                         GM_ADDR value, 
                                                         GM_ADDR z);
#endif

int32_t main(int32_t argc, char *argv[])
{
    uint32_t blockDim = 8;
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    size_t maskByteSize = 8 * 2048 * sizeof(bool);
    size_t valueByteSize = sizeof(uint16_t);
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);

#ifdef ASCENDC_CPU_DEBUG
    uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputByteSize);
    uint8_t *mask = (uint8_t *)AscendC::GmAlloc(maskByteSize);
    uint8_t *value = (uint8_t *)AscendC::GmAlloc(valueByteSize);
    uint8_t *z = (uint8_t *)AscendC::GmAlloc(outputByteSize);

    ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
    ReadFile("./input/mask.bin", maskByteSize, mask, maskByteSize);
    ReadFile("./input/fill_value.bin", valueByteSize, value, valueByteSize);

    AscendC::SetKernelMode(KernelMode::AIV_MODE);
    ICPU_RUN_KF(mask_fill_custom, blockDim, x, mask, value, z);

    WriteFile("./output/output_z.bin", z, outputByteSize);

    AscendC::GmFree((void *)x);
    AscendC::GmFree((void *)mask);
    AscendC::GmFree((void *)value);
    AscendC::GmFree((void *)z);
#else
    CHECK_ACL(aclInit(nullptr));
    int32_t deviceId = 0;
    CHECK_ACL(aclrtSetDevice(deviceId));
    aclrtStream stream = nullptr;
    CHECK_ACL(aclrtCreateStream(&stream));
    uint8_t *xHost, *maskHost, *valueHost, *zHost;
    uint8_t *xDevice, *maskDevice, *valueDevice, *zDevice;

    CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputByteSize));
    CHECK_ACL(aclrtMallocHost((void **)(&maskHost), maskByteSize));
    CHECK_ACL(aclrtMallocHost((void **)(&valueHost), valueByteSize));
    CHECK_ACL(aclrtMallocHost((void **)(&zHost), outputByteSize));
    CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void **)&maskDevice, maskByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void **)&valueDevice, valueByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void **)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

    ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
    ReadFile("./input/mask.bin", maskByteSize, maskHost, maskByteSize);
    ReadFile("./input/fill_value.bin", valueByteSize, valueHost, valueByteSize);

    CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
    CHECK_ACL(aclrtMemcpy(maskDevice, maskByteSize, maskHost, maskByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
    CHECK_ACL(aclrtMemcpy(valueDevice, valueByteSize, valueHost, valueByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
    mask_fill_custom_do(blockDim, stream, xDevice, maskDevice, valueDevice, zDevice);
    CHECK_ACL(aclrtSynchronizeStream(stream));

    CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
    WriteFile("./output/output_z.bin", zHost, outputByteSize);

    CHECK_ACL(aclrtFree(xDevice));
    CHECK_ACL(aclrtFree(maskDevice));
    CHECK_ACL(aclrtFree(valueDevice));
    CHECK_ACL(aclrtFree(zDevice));
    CHECK_ACL(aclrtFreeHost(xHost));
    CHECK_ACL(aclrtFreeHost(maskHost));
    CHECK_ACL(aclrtFreeHost(valueHost));
    CHECK_ACL(aclrtFreeHost(zHost));
    CHECK_ACL(aclrtDestroyStream(stream));
    CHECK_ACL(aclrtResetDevice(deviceId));
    CHECK_ACL(aclFinalize());
#endif
    return 0;
}
